test(BACKEND-ROCM): the block-size contract is enforced at its production call site - #1273

…-size reachability test

Picks up `e2a9e035d` (#1273). No conflict: it appends a case to
`tests/vllm/v1/worker/test_runner.cpp`, which this row does not touch. The
merge is taken so the second-review repair runs its gate on the tree that will
actually land.

Row: SPEC-DFLASH2-W6

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…-pin vLLM oracle, and the speed axis is owed (#1314) (#1637)

DFlash2 (`DFlash2DraftModel`) now drafts on this engine, token-exact
against a
beyond-pin vLLM oracle. W6 is the final wave of the six-wave port and
the one
that reads the oracle: W1-W5 landed the config surface, the grouped
dynamic
convolution, the candidate selector, the pair-emitting top-k, the path
walk and
the GGUF arm.

WHAT THE GATE MEASURED. G2 compares our draft blocks against a vLLM
capture on
the published `z-lab/Qwen3.8-27B-DFlash2` checkpoint: 4 of 4 prompts
token-exact,
45 of 47 draft blocks byte-identical. G3 reads acceptance per prompt --
49/54/54/52 accepted over 15/10/10/12 blocks -- identical on both
engines. The
two divergent blocks reject the divergent slot on BOTH sides, which is
why `len`
stays in sync and why G3 is recorded as a COROLLARY of G2 rather than an
independent instrument: on this capture it is arithmetically entailed.

WHAT THE ORACLE DID. Running vLLM against itself across its own
attention
backends, the outputs diverge on 1 of 4 prompts and acceptance moves
0.597 vs
0.657. That is upstream disagreeing with itself, recorded because it
bounds what
any single-backend comparison can claim.

THE PIN IS A DATED EXCEPTION. vllm-project/vllm#52816 MERGED 2026-08-21
at
05:27:22Z (merge commit `b389ac29`, head `3406ec1d`). The gate stays
pinned at
`66e5414c` because that is the wheel that executed; re-labelling a
capture to a
head it never ran would be a false pin. #1561 owns moving it.

The `#1538` index row was SUPERSEDED, not corrected:
`.agents/issue-index.md` is
append-only and a landed row is never edited, so the first repair wave
appended
#1561 rather than editing #1538. The spec's `## Upstream chain` table is
the
authority on what that row's superseded text says.

WHAT IS OWED, NAMED RATHER THAN LEFT TO BE FOUND.

- The SPEED axis is NOT TAKEN. No idle-host A/B was reached, our draft
is off
the paged CUDA-graph fast path, and wall clocks were
CIFS-load-dominated. This
  is an open gap, not a ceiling.
- O24: the `at_end` fixture is committed but its `<` to `<=` mutation
was never
run in-suite. An out-of-suite compilation of `ReconstructAcceptance`
shows the
case DOES discriminate (`verified` 1 vs 2), which removes the risk that
the
  fixture proves nothing -- it does NOT discharge the owed run.
- O25 / #1607: the DFlash2 startup notice prints TWICE per load.
`CheckDflash2DraftArm`
is reached from `FromModelDir` and again through the engine
constructor's
speculative-config resolution, and carries no once-flag. No weights are
loaded
  twice; the classification and its paragraph re-run.
- #1562: the capture harness exists only as prose, so the FLASH_ATTN
golden's
backend label is a post-hoc relabel that cannot be re-derived from this
tree.

INHERITED RED. `test_runner`'s "refuses a non-multiple-of-16 block size"
case
fails on this tree and on pristine `origin/main` for x86 CPU builds. It
is #1608,
introduced by #1273 and owned by `BACKEND-ROCM`: `build-test-cpu` was
green at
`483cd319` and red at `e2a9e035`, while `build-test-cpu-arm64` passes at
both.
Not caused here and not repaired here. `test_ltx2_video` (#1559) is
load-sensitive
and passed on the sibling branch at the same pin.

Full ctest on the merged tree: 586 tests, 581 passed, 3 skipped, 2
failed, both
named above.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: ClaudeCode:claude-opus-5 [Claude Code]

---------

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>

…sed at the loader because nothing resolves a chain yet (#1522) (#1638)

`--speculative-config` gains an OPTIONAL `vllm_cpp.drafter_chain` field:
a
preference-ordered list of speculators. W1 lands the field and a
loader-level
refusal; nothing resolves a chain yet, so a specified chain is refused
by name.
W2 (attribution) and W3 (resolution) follow.

ADDITIVE, MEASURED AGAINST THE PIN. `vllm_cpp` and `drafter_chain`
appear nowhere
in vLLM's `SpeculativeConfig` at the parity pin `555967922`, read from a
local
clone at that exact SHA. Exactly one name is added and it is nested,
following
the precedent `src/vllm/config/weight_residency.cpp` already sets. No
vLLM-owned
field is renamed, re-typed, shadowed, or has its parse or validation
semantics
changed.

INERT WHEN ABSENT, KEY FOR KEY. A differential harness compiled the
config parser
at the merge base and at this head and ran an identical 203-document
corpus of
chain-free `--speculative-config` documents through both: all 11 method
spellings,
all 26 upstream-unimplemented keys with and without values, every
honoured key
crossed with 8 malformed value types, unknown-key spellings including a
bare
top-level `drafter_chain` and a mis-cased `VLLM_CPP`. Identical
accept/reject
verdict, identical resolved fields, identical error ordering on every
one.

The one delta, stated rather than waived: 152 refusal messages differ by
the
appended accepted-key list, because that list must name the new key or
it stops
closing the search. "Key for key" is true; "byte-identical" would not
be, and an
earlier draft of this row claimed it.

Structurally, `cfg.drafter_chain` has exactly one writer, inside `if
(has_chain)`
where `has_chain` is `doc.contains("vllm_cpp")`, so the vector is
provably empty
for any document not naming that key and both loader guards are provably
false.

REACHED, PROVEN BY DELETION. Deleting the `ResolveSpecConfig` refusal
reddens 4
cases; deleting the `FromModelDir` guard reddens 3. Both enter through
`include/vllm.h` (`vllm_engine_load`), not through the type. The refusal
fires
before any weight I/O -- the tests point at a nonexistent model path, so
a guard
placed one line too late answers "model path is not a directory" and the
case
reddens.

A REAL vLLM KEY WAS BEING CALLED A TYPO, FIXED IN FLOW (#1598). Inside a
chain
entry, `draft_sample_method` and `rejection_sample_method` fell through
to the
unknown-key message, though both are genuine `SpeculativeConfig` fields
honoured
at the top level. They now share one named value-gated set read by both
admissions, and an entry refusal says the engine honours the key and to
spell it
at the top level. Red-first: 4 failing assertions before the fix.

Gates: `test_speculative_drafter_chain` 13 cases / 149 assertions;
`test_drafter_chain_reach` 7 cases / 78 assertions. 6 of 6 mutations
detected,
each with match count, diff stat and compile rc recorded; two of them
confirm the
new tests did not disarm the original reachability proofs.

OWED, and filed rather than left to be found: #1599 (a non-string
`model` on an
entry is dropped, then reported missing -- inherited from the
pre-existing
top-level behaviour), #1600 (a mis-cased `vllm_cpp` reports the method
requirement instead of naming the key), #1601 (llama.cpp cited by line
with no
revision against oracle pin `10bf611e`).

INHERITED RED. `test_runner`'s "refuses a non-multiple-of-16 block size"
fails
here and on pristine `origin/main` for x86 CPU builds: #1608, introduced
by #1273,
owned by `BACKEND-ROCM`. Not caused here, not repaired here. Full ctest
on the
merged tree: 587 tests, 583 passed, 3 skipped, 1 failed -- that one.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: ClaudeCode:claude-opus-5 [Claude Code]

---------

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>

…ow says why, and AttentionDenseFlash advertises the head_dim it can launch (#1544) (#1578)

Two additive changes from #1544's `## Owed`. Neither moves a single
existing
caller's numerics, and that constraint shaped the whole design.

## The naive rung stops being a silent default

`vt::Attention` resolves `OpId::kAttention` straight to the
correctness-grade
kernel, and nothing in the tree ever routes it up: the rung is whichever
C++
function name the author typed. That is deliberate for six of the nine
call
sites, and invisible to everyone else, which is how one LTX-2.5 DiT
forward came
to cost 47.84 s. A token gate cannot see the difference by construction
— every
rung is bit-identical or inside the bf16 envelope, so the goldens pass
either
way.

A selector that auto-routes was rejected, and not on taste. Three of the
six
sites are reference arms a gate compares against (`nemotron_h.cpp`,
`nemotron_h_device.cpp`, `qwen3_5.cpp`), two are the `VT_*_EAGER` rungs
of a
same-binary A/B (`whisper_audio.cpp`, `qwen3_vl_vision.cpp`), and one is
a
measured-negative device path behind `VT_KIMI_DEVICE_MLA`. Rerouting any
of them
changes what the reference computes, which deletes the comparison the
gate
performs rather than fixing anything — the "widen the assertion until
the gate
passes" failure AGENTS.md names. `kAttention` and `vt::Attention` are
untouched
here.

`scripts/check-attention-rung-consistency.py` requires the CHOICE to be
recorded
instead: a `// VT-ATTN-NAIVE:` reason on the call line or within 20
lines above
it. The six deliberate sites now carry one, and an author who never
heard of the
fast rungs gets a red instead of a silent 500x. The scan runs over
`checker_text.normalize_source`, so a commented-out, `#if 0`-ed or
`if (false)`-ed call is a deletion to it exactly as it is to nvcc, and
the
reported `file:line` still describes the original file.

The record is per-site and in-file, so an ordinary change writes no
shared
record at all. `scripts/attention-rung-allowlist.txt` holds only the
three stems
whose naive call another row is currently deleting —
`muse_glimmer_vision`
(#1545) and the two LTX-2.5 files — because editing the very lines those
changes
replace would conflict for no gain. A stale entry there is reported and
is not
fatal, so the removing row owes this file nothing.

## `AttentionDenseFlash` advertises the head_dim it can launch

It claimed `head_dim <= 256` while asking the driver for
`2*kFlashBc*d*sizeof(Tin)` bytes of dynamic shared memory, with no
`cudaFuncSetAttribute` anywhere in `src/vt/cuda/`. The default 48 KiB
cap made
the real ceiling 192 in bf16 and 96 in f32, so Kimi at 192 f32 or
Qwen3.5 at 256
would have received a bare launch error from the `cudaGetLastError` at
the
bottom of the launcher, naming nothing they could do instead.

The bound now lives in `include/vt/ops.h` as
`AttentionDenseFlashSmemBytes` and
`AttentionDenseFlashMaxHeadDim` — pure host arithmetic, so a box with no
GPU can
execute it — tied to the kernel by two `static_assert`s on the tile
width and
the register blocking. The launcher refuses above it naming
`vt::AttentionDenseFast`, which uses no shared memory and does serve
those
widths.

Narrowing beats opting in to a larger cap here, and that is now a
measurement
rather than a preference. **GB10's queried opt-in ceiling is 101,376
bytes**
(measured during #1557's review), while head_dim 256 in f32 wants
131,072. So
`cudaFuncSetAttribute` cannot make the widest advertised width true on
the part
this project gates on — the raise buys nothing at the width that
motivated it,
and a caller there would have gone on falling back silently without ever
launching. The bound is INCLUSIVE, which matters in one direction:
head_dim 192
in bf16 lands exactly on 49152 and launches today, so an exclusive bound
would
refuse work that currently runs. Opting in stays available later as a
widening
for bf16 above 192, owned by nobody today.

This mirrors vLLM's own polarity rather than inventing one:
`vllm/model_executor/models/vision.py:99` selects an encoder backend by
shape,
and `vllm/v1/attention/backend.py:155-163` consults `supports_head_size`
BEFORE
dispatch instead of discovering the domain by launching. Both read at
the pinned
oracle `555967922`.

## Evidence

RED first, on the unmodified tree: the checker reported all six
deliberate sites
at the exact lines #1544 names (`kimi_linear_device.cpp:598`,
`nemotron_h.cpp:671`, `nemotron_h_device.cpp:330`, `qwen3_5.cpp:5279`,
`qwen3_vl_vision.cpp:527`, `whisper_audio.cpp:324`) and correctly
excluded the
three allowlisted ones. GREEN after the markers: 9 sites, 6 marked, 3
allowlisted.

`tests/scripts/test_check_attention_rung_consistency.py` 34/34,
including six
mutations that must go RED — a new unmarked model, a new unmarked call
in a
HEADER, a deleted marker, a second unmarked call inside an
already-marked file, a
stub reason, and a widened regex that would swallow the fast rungs. It
also pins
that the scanned population is not empty, which is the guard against the
way a
structural checker usually goes green: by matching nothing at all, and
that every
allowlisted stem names a model source that exists, which is what catches
a typo.

`tests/vt/test_ops_attention.cpp` gains the head_dim contract cases: the
tile
arithmetic at both element sizes, both honest bounds, that 256 is
outside both,
and the inclusive edge in both directions. 11 cases / 39 assertions,
SUCCESS.

MUTATED, because a green suite over new arithmetic proves only that the
arithmetic agrees with itself. Making `AttentionDenseFlashMaxHeadDim`
return
`kAttentionDenseMaxHeadDim` — the exact contract this change repairs —
turns the
new case RED on 6 of its assertions, each printing the wrong value it
now
carries (`256 == 192`, `256 == 96`, `65536 <= 49152`, `131072 <=
49152`), so the
mutation demonstrably applied and demonstrably compiled.
`include/vt/ops.h`
restored and verified by sha256 against its pre-mutation snapshot;
rebuilt; 39/39
green again.

Proof that no caller's numerics moved: the checker executes no model
code; the
head_dim guard fires only where the launch already failed; no marker
changes a
statement; and `git diff` touches no kernel arithmetic, no dtype and no
default.

## What the fresh review changed

Six findings, repaired here. None of them moves a kernel's arithmetic
either.

`pr-size` was RED, and this branch caused it. A checker created inside
the range
has no BASE version for the red-before half of the evidence run, so it
has to
register the disabled stub its own suite must reject; about twenty
checkers do,
and this one did not, so the gate could not classify the change at all.
Measured
rather than asserted: under the stub every case goes red — re-measured
after the
repairs below, `FAILED (errors=34)` — because the suite loads the
checker as a
module and every case calls into it. `agent-preflight.sh` does
not run `check-pr-size`, which is why a local green said nothing about
it.

The second `static_assert` beside the kernel was a tautology. It read
`8 * 32 == kAttentionDenseMaxHeadDim` while the real `kMaxPerLane` was a
function-local `constexpr` inside the kernel body, invisible at file
scope, so
setting that local to 4 — precisely the drift the message claims to
catch — left
the assert reading `256 == 256`. The register blocking is now
`kFlashMaxPerLane`
at file scope; the kernel's register arrays and unrolled loops read it,
and so
does the assert. The same mutation now reads `128 == 256` and fails to
compile.
There is no nvcc on this box, so the tie was measured by extracting that
constant
block from `cuda_ops.cu` VERBATIM and compiling it against the shipped
`include/vt/ops.h` under `g++ -fsyntax-only`: clean before, `static
assertion
failed` after, `cuda_ops.cu` restored and verified by sha256. The first
assert
(`kFlashBc == kAttentionDenseFlashTileCols`) was already a real tie and
is
untouched.

Two comments claimed more than the code delivers. The launcher said its
guard and
its shared-memory request came from "the SAME function … cannot
disagree"; they
are two functions, and `AttentionDenseFlashMaxHeadDim` re-derives the
division
instead of inverting `AttentionDenseFlashSmemBytes`. The guarantee holds
and is
tested: mutating the `2 *` in `SmemBytes` to `3 *` reds 9 assertions of
the
shipped contract case, both inclusive-edge checks among them, while
`MaxHeadDim(2) == 192` stays green — which is the re-derivation made
visible. The
comment now describes that. The `AttentionDenseFa2` fall-through comment
promised
"the best available kernel for their shape rather than a hard refusal",
which
stopped being true for an over-cap head_dim the moment this branch added
the
refusal; it now names the domain and records that every caller today is
far
inside it (max head_dim 80).

The checker claimed "the population is what makes a green meaningful"
and named
no limits. Four spellings reach the same kernel undetected — a `using`
declaration, a namespace alias, a `#define`, and a call through
`&vt::Attention`
— each verified during review to leave the checker green with a live
unmarked
call. None exists in this tree, and widening the regex would make every
fast rung
a site, which is D1's rejected failure mode again; closing it needs a
compiler-side population, not a longer pattern. The docstring and spec
D6 state
the bound, so a green reads as "no unmarked `vt::Attention(` call" and
never as
"no model is naive".

The OK line reported total and marked sites but never the number a
reader needs:
sites carrying no reason that pass only because their stem is
allowlisted. It is
not `sites - marked`, since a marked call inside an allowlisted file
counts in
`marked`. Two cases now pin the line; dropping the count from it reds
them.

Two records were wrong. `scripts/attention-rung-allowlist.txt` told a
removing row
to delete its stem without saying that
`test_allowlist_holds_only_the_in_flight_stems`
pins the set in another file and reds on the deletion; the allowlist
header, the
checker docstring and spec D7 now say so. The kernel-matrix cell stored
this
suite's case count — a measurement of one file inside another, which
AGENTS.md
names as a drift lock — so the count is gone rather than corrected.

## Two drift locks in the new suite, both repaired here

Found while landing this change against #1579, by checking the
interaction
instead of assuming the two pull requests were independent. Each was
green on its
own; `main` went red only once both landed, which is why nothing on
either branch
caught it.

**The first was the population floor.**
`test_the_population_is_not_empty`
asserted the scanned population was `>= 9` against a tree of EXACTLY 9
sites. Its
own name says "is not empty" and its assertion pinned a count: the name
was right.
A raw total is a measurement of the model tree stored in a test file,
which
AGENTS.md `## Records` forbids, and it reds on any row that legitimately
REMOVES a
naive call — every stem on `scripts/attention-rung-allowlist.txt`, which
is to say
the rows that allowlist exists to unblock. It runs in the required
`agent-record`
job, so #1545 alone would have turned `main` red.

**The second was `assertGreater(excused, 0)`** in
`test_the_ok_line_reports_the_excused_sites`. Same shape, one case down:
`excused`
counts unmarked calls in allowlisted files, so it reaches 0 when the
LAST stem is
cleaned up, redding the case while the checker is green at rc=0. It does
not fire
for any of the three rows individually, so it would have sat latent
until the
LTX-2.5 reroute tripped it.

Both are repaired here rather than deferred, because none of this has
landed:
`scripts/check-attention-rung-consistency.py`, its allowlist and its
suite are all
CREATED by this pull request, so correcting a defective assertion in
them is
repairing the change, not amending a gate that `## Changing the rules or
a
checker` governs.

The floor is now genuine non-emptiness (`>= 1`). One new case asserts
every
allowlisted stem NAMES AN EXISTING model source — keyed on file
existence and
deliberately never on scan membership, because a stem stops having a
call site the
moment its removing row lands, which is the state the allowlist is built
to
survive and which `stale_allowlist_entries` already promises in its own
docstring
("Reported, never fatal"); asserting scan membership would have rebuilt
the
identical lock one line over. The `excused > 0` floor is gone, its
shipped-tree
half kept because it RE-DERIVES the count instead of pinning it and so
holds at 0
as well as at 3, and the coverage it was standing in for moved onto two
cases
driven over a constructed scan and a temporary allowlist, which the
model tree
cannot switch off.

RED first, each mutation proven applied and restored by sha256. Stubbing
`scan_models` to `{}` reds the floor at `0 not greater than or equal to
1`, and
independently so does renaming the checker's regex. A typo'd
`muse_glimmer_vison`
entry reds the new case naming that stem while the CHECKER stays green —
which is
the point, since a bogus stem is reported only as STALE and never fails.
For the
second lock the control is sharper: with the checker's excused count
broken AND
the tree in its end state, the retained shipped-tree case goes GREEN
over the
defective checker and only the constructed case catches it, so the
replacement is
real coverage rather than a deletion in disguise.

Composed green, measured on this head with #1579's
`muse_glimmer_vision.cpp`
copied in: with the allowlist stem left in place the suite is `Ran 34
tests ... OK`
and the checker prints `STALE (not a failure)` at rc=0, where before the
repair it
read `8 not greater than or equal to 9`. With the stem also deleted,
only
`test_allowlist_holds_only_the_in_flight_stems` reds, which is the
by-design pin
on the set, and updating that set in the same change returns it to green
— so both
removal routes now have one.

A fresh review of the repair returned one finding, repaired here, and
repairing
it turned up a second of the same kind. The new case's comment claimed
it also
caught a checker printing `sites - marked`. It does not, and no
constructed scan
could make it: `main()` reaches the OK line only when `drift_sites` is
empty, and
then every unmarked site is excused, so the two quantities coincide
identically —
1000 reachable green states enumerated, 0 where they differ.
Substituting
`sites - marked` into the checker leaves all 34 cases green. Nearby,
`test_the_excused_count_is_not_sites_minus_marked` claimed that dropping
the
allowlisted file's marked site made the two diverge; it does not,
because that
lowers `sites` and `marked` together. Both comments now state what is
actually
pinned and why the rest cannot be, which is the same "a comment claimed
more than
the code delivers" class this branch already repaired twice. The fix is
comment-only, and that is proven rather than asserted: `ast.dump` is
byte-identical
across the change, so no assertion, fixture or docstring moved.

Five rounds of comment repair were needed, because each of the first
three
removed a false claim by writing a NEW causal explanation that the next
fresh
review then measured false. From the third round on the rule was to
DELETE rather
than re-explain, and to run every clause left standing. Comment lines go
down, not
up. `ast.dump` is byte-identical across rounds three and four, so no
assertion or
fixture moved; round five changes exactly one `assertTrue` MESSAGE
string, with the
assertion condition's own `ast.dump` hash shown identical either side.

Two findings from those reviews were REFUTED by measurement rather than
applied.
`test_widening_the_regex_to_the_fast_rungs_is_visible` was reported dead
because it
survives a `drift_sites` break and a `\b` removal — but neither is the
widening it
names, and mutating `_NAIVE_CALL` to `\bvt::Attention\w*\s*\(` reds it
along with
five others. The spec's "six mutations that must go RED" was reported as
five; there
are six, and each reds under the mutation it names. Applying either
would have
renamed a working test and made an accurate record false.

**Knowingly shipped, and recorded rather than hidden.** Three comments
in
`scripts/check-attention-rung-consistency.py` (`:58-61`, `:93-96`,
`:252-255` —
anchors measured, not estimated) are MEASURED FALSE and are NOT repaired
here. All
three are now enumerated in the test file beside the case that pins the
real
behaviour, so a reader of the suite can find every one. They claim the
`\b` in
`r"\bvt::Attention\s*\("` is what stops the pattern matching the fast
rungs. It
is not: the trailing `\(` does that, `vt::AttentionDenseFlash(` matches
with
neither, and the two patterns differ on exactly the 63 identifier
characters —
`xyvt::Attention(` alone. With the `\b` removed the suite stays green at
`Ran 34 tests ... OK`, so the companion claim that the widening "is
caught in this
suite" is false too. The equivalent claims in the TEST file ARE
repaired, because
`tests/scripts/test_*.py` is not a governance checker. **The tree
therefore
contradicts itself between the test and the checker beside it, and that
is
deliberate**: `scripts/check-pr-size.py:170` classifies every
`scripts/check-*.py` as a `governance_checker` and demands executable
red-before
evidence, which a comment-only diff cannot produce. Attempting the
repair returns
`ERROR: BASE checker stayed green ... changed test is not semantic
evidence` at
rc=1. The prepared patch was deliberately not committed rather than land
a red
`pr-size`.

#1629 records both drift locks. #1631 records the comment freeze, and it
is not
one file: the pattern covers all 43 `scripts/check-*.py` checkers plus
the `.sh`
ones, so a false comment in any checker in this repository cannot be
corrected on
its own today.

Both are linked in the three places AGENTS.md requires rather than in
this body
alone: `.agents/issue-index.md` gains one appended row each, and `##
Owed` in
`.agents/specs/attention-rung-visibility.md` records #1629 as DISCHARGED
HERE and
#1631 as owed with the reason it cannot be. No gate would have caught
their
absence, because `check-agent-record.py` counts index rows that name no
owner and
there was no row at all.

## Owed, and named rather than skipped

Nothing on a CPU-only box executes `LaunchAttentionDenseFlash`, so the
pure-arithmetic cases stay green over a launcher that lost the bound.
The
on-device refusal case emits a loud PENDING message and returns, and
#1573 owns
running it plus the reachability mutation that proves the case reaches
the
guard. No lease was taken: `dgx:gpu0` was unavailable for the whole
branch. The
101,376-byte GB10 ceiling above does NOT discharge #1573 — it bounds
what an
opt-in could buy and says nothing about whether the launcher's refusal
executes.

Owed and filed rather than left to be discovered: **#1629**. The new
`test_the_population_is_not_empty` case asserts the scanned
`vt::Attention`
population is `>= 9`, and the shipped tree has exactly 9 sites, so the
floor has
ZERO headroom and any row that deletes a naive call reds it. That is
every stem
on `scripts/attention-rung-allowlist.txt` — `muse_glimmer_vision`
(#1545) and the
two LTX-2.5 files — which is to say the rows the allowlist exists to
unblock.
Measured while landing this change, with #1579's
`muse_glimmer_vision.cpp` copied
onto this head: leaving the stem reds that case at `8 not greater than
or equal
to 9`, and deleting the stem reds it AND
`test_allowlist_holds_only_the_in_flight_stems`, so a removing row has
no green
path. The tree was restored byte-for-byte after each and the suite
returns to
`31 tests ... OK`. This is the same drift-lock shape the kernel-matrix
cell above
was corrected for, retained one file away, and it is NOT repaired here
because
changing the floor changes what the gate accepts and AGENTS.md routes
that to its
own row, spec and red-before evidence. #1629 carries the evidence and
two
candidate directions, and #1579 is held on it.

Inherited red, not introduced here: `test_cpu_x86_llamacpp_floor` fails
in
`agent-preflight.sh` on this box. It is the known load-dependent case of
#618 —
at high loadavg the harness exits `NO_QUIET_WINDOW (4)` where the case
expects
`GIVING_UP (2)` — and both the case and
`scripts/cpu-x86-llamacpp-floor.sh` are
byte-identical to `origin/main` on this branch, which is how it was
established
as inherited rather than assumed to be.

FIVE CI jobs are red on the head, every one of them inherited from main
with a
named owner, and each was verified per-job against a main baseline
rather than
asserted. `windows-msvc-cpu` and `windows-msvc-vulkan` are the standing
PR-only
red (#584, #965). `build-test-cpu` and both `sanitize-cpu` arms fail on
ONE
shared doctest case, `test_runner.cpp:1557`; that case landed on main in
`e2a9e035d` (#1273) and is owned by #1608 and #1602. Inheritance was
established
by comparing the failing ASSERTION and not the job name: the scheduled
main
baseline at `e2a9e035d` fails the identical
`CHECK_THROWS_WITH_AS( make_runner(), "Block size must be a multiple of
16", std::invalid_argument )`
at the same `test_runner.cpp:1557`, with the same "No valid attention
backend for
device type 0" text, in all three jobs, and with zero sanitizer findings
in either
sanitized arm. `test_runner.cpp` is not touched by this branch.

`build-newest-gcc` was the sixth red when this body was first written
and is
GREEN here. #1581 and #1618 landed the `::getpid` repair on main, and
this branch
was re-merged onto `2e7f3bee7` to pick it up, so the job now compiles
and reports
on this change. Every job that can see this change is green, including
`pr-size`
and `agent-record`, which is the job that runs this row's new checker.

Closes #1544.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]

---------

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>

… handing out an address-keyed cache (#1415)

Row `A2-Q2b`, issue
[#810](#810), spec
`.agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md`.

Measured first, built second. The row's premise was that NemotronH
decode re-expands quantized weights on the host and that `lm_head` is a
large share of it. That was arithmetic, and the brief that dispatched
this row said so. It is now a count.

**This description was rewritten after a fresh review returned
FINDINGS.** The measurement and the seam design below survived review
unchanged and were verified independently. Three claims that were in the
previous description did not, and they are corrected in place rather
than quietly dropped: the new test had never compiled, the routing
checker was said to hold a branch it does not hold, and the Marlin arm
was said to be unbuildable here when it builds in seconds.

**A scoped re-review then found no correctness defect in the shipped
code and verified all seven repairs, and a final round corrected what it
did find — all of it evidence, records and process.** The evidence
tables named a tree they were not measured on and now carry an `overlay`
column with the un-overlaid baseline stated as its own row; the Marlin
object figure is re-measured and anchored to a SHA; the shared
process-static this row inherits is named; the #1421 index row's counts
are corrected while the pull request is still unmerged and the row still
editable; and `§4` of the spec stopped telling the next agent to take
`flock` on a fleet device. No product code changed in that round.

## What the measurement says

Tallied at `NemotronHOwned::DenseBf16` (`nemotron_h.cpp:395`), the
single seam every host dequant in this architecture passes through, on
the real 21 GB checkpoint (revision `29f2d174`) through the production
ABI driver `examples/nemotron_h_gen`. One decode step, T=1, `top_k` 6,
23 MoE layers, 369 dequant calls:

| group | shape | calls | elements | % |
|---|---|---|---|---|
| routed expert `up_proj` | `[1856, 2688]` | 138 | 688 472 064 | 22.36%
|
| routed expert `down_proj` | `[2688, 1856]` | 138 | 688 472 064 |
22.36% |
| shared expert `down_proj` | `[2688, 3712]` | 23 | 229 490 688 | 7.45%
|
| shared expert `up_proj` | `[3712, 2688]` | 23 | 229 490 688 | 7.45% |
| `lm_head` | `[131072, 2688]` | 1 | 352 321 536 | 11.44% |
| mamba `out_proj` (FP8) | `[2688, 4096]` | 23 | 253 231 104 | 8.23% |
| mamba `in_proj` (FP8) | `[10304, 2688]` | 23 | 637 034 496 | 20.69% |
| **TOTAL** | | **369** | **3 078 512 640** | **100%** |

`138 == 6 x 23` exactly, which is what identifies this as the decode
shape rather than a prefill aggregate.

Three findings, and the first two are corrections:

1. The dispatching estimate put `lm_head` at 131072 x 4096 = 537e6
elements and ~43% of the population. `hidden_size` is 2688. It is 352
321 536 elements and 28.35%.
2. The widely quoted "1.24e9 elements / 2.49 GB per token" names a
REGIME, not a total. It is not the host arm's 3 078 512 640. It is
`mamba + lm_head` = 1 242 587 136 = 2.485 GB, the residue after A2-Q2a
moved the MoE arm to the device, matching to four significant figures.
3. `lm_head` is the last one. Against the A2-D1 discriminator on
`dgx:gpu0` (device mamba ON 1.554 s/token and 108.2x vs vLLM, OFF 10.319
s/token and 718.1x), the mamba arm is worth 6.64x and `lm_head` is on
the host in every leg. Once the mamba arm lands, `lm_head` is 352 321
536 of 352 321 536, i.e. everything that remains. It is also the largest
single re-expansion in the model by 12.7x, so its 704.6 MB transient
bf16 buffer is the one that matters on a unified-memory box that reboots
rather than OOM-kills.

So the estimate's share was wrong, its direction was right, and the case
is stronger after the discriminator than before it.

## The residency decision, taken explicitly

The spec's `## 5. Owed` required this row to choose rather than default:
"A2-Q2a routed around it by never calling either
`MarlinDenseResidentFor`; `lm_head` must do the same or say why not."
That accessor keys its repack cache on the weight's ADDRESS
([#984](#984)), and NemotronH
is exactly the second-consumer condition an address key cannot survive.

Rather than fork the seam as `qwen3_5.cpp` did, or inherit the defect as
routing through it plainly would, the seam is EXTENDED:
`MatmulNvfp4W4A16D` and `MatmulNvfp4MarlinD` accept a caller-owned
`MarlinDenseResident*`, defaulting to the existing cache so every
current caller is byte-for-byte unchanged. NemotronH owns its resident
in a `ResidentSlot` on the weights, which is the property
[#237](#237) added
`ResidentSlot` for and the same call A2-Q2a made for the MoE arena.
[#984](#984) is left exactly as
it was for every other caller, where its own two-engine red-before
belongs.

## What moved

`NemotronHPagedForward`, the production path, now returns a
device-resident `ForwardLogits` on an NVFP4 checkpoint, which is every
released one, and downloads `final_normed` only when a consumer needs
it.

`NemotronHDeviceForward` deliberately keeps the host projection. It is
the operand the numeric gate compares against, it is what makes A2-R's
token gate attributable, and it has no production caller. The spec
required this row to disclose where that property ends: it ends in the
paged forward.

## The allowlist entry is narrowed, not removed

Its one clause is met. But `check-runner-routing-consistency.py`
resolves a cross-TU delegate only for the `Class::ForwardDevice` shape,
and NemotronH's device forward is a free function in another TU, so the
model still classifies HOST although `NemotronHPagedForward` assigns
both `fl.device_tensor` and `fl.device_storage`. Filed as
[#1410](#1410) and left to its
own row: it changes checker semantics, which `AGENTS.md` routes away
from an in-flow fix, and widening a classification to turn a red gate
green is precisely the move that section slows down.

## The review repair

### The new test had never compiled, so the red it claimed could not
have existed

`tests/vllm/models/test_nemotron_h_moe_device.cpp` used
`NemotronHHostWeights` unqualified and never declared it. Measured on
the same command:

| tree | `compile_rc` | errors |
|---|---|---|
| PR head `29b1128e`, this file | 1 | 9 |
| merge base, same file | 0 | 0 |
| repaired, default arm | 0 | 0 |
| repaired, `-DVT_MARLIN_NVFP4=1` | 0 | 0 |

The consequence is larger than the declaration. The file had never
compiled on any platform, so **the red-first result this row claimed for
its synthetic numeric gate did not exist and could not have existed**.
That gate is CUDA-only and still has never executed; the spec now
records `NEVER RUN` rather than restating a red nobody saw. The test
target is deliberately registered with no CUDA guard, and that is what
surfaced this — a case that skips at run time still has to parse on
every CPU build.

### The production source asserted a protection that does not exist

A comment at the device branch claimed the allowlist entry was removed
and that "the routing checker, not a comment, is what now holds this
branch in place: put the host projection back and it goes red by name".
All three claims were false, and they contradicted this description and
the allowlist file in the same commit. Reproduced here on the repaired
tree: baseline `check-runner-routing-consistency.py` exits 0 with "3
host-logits off-framework (3 allowlisted)"; after deleting the whole `if
(DeviceLmHeadEligible(...)) { ... return fl; }` block (`git diff` moves
by 1 insertion and 14 deletions, so it applied) the checker exits 0 with
**byte-identical output**; the file was restored to an identical sha256.

**Nothing automated holds that branch.** The checker is not widened to
make it — that changes checker semantics and
[#1410](#1410) owns it with its
own red-before. The comment now says what is true, the allowlist entry's
own text is corrected the same way, and the reachability deletion
mutation is recorded PENDING a CUDA window rather than claimed. Every
case that can enter the device branch is `MarlinW4A16Selects`-gated and
therefore CUDA-only, so that mutation cannot run on a CPU box.

### A silent dequant fallback that no token gate can see

`DeviceLmHeadEligible` restated the shared dispatcher's three selection
clauses and dropped `MarlinW4A16Enabled()`. Under an explicit
`VT_NVFP4_MARLIN=0` the model said eligible while `MatmulNvfp4W4A16D`
took its naive redundant-dequant arm — the same logits, while
re-uploading the whole `[131072, 2688]` operand on **every decode
step**, because `LmHeadNvfp4View` hands out a stack temporary that
`ResidentNvfp4`'s weight-keyed cache can never hit.

The repair is structural rather than a patched clause. The three clauses
now live once, in `dense_nvfp4::MarlinW4A16Selects`, and both the
dispatcher and the model call it, so they cannot drift. `DeviceLmHeadD`
additionally refuses BY NAME if that same predicate is false on the
operand it is about to hand over, which catches an eligibility answer
taken against a different queue or dtype. That is deliberately the
predicate and not the seam's `fallback_gemms` counter:
`MutableW4A16Stats()` is a plain non-atomic process-wide static, so a
counter window in production would refuse a correct run whenever
anything else took a fallback GEMM concurrently, and a false refusal is
worse than the silence it replaces. The counter is the right instrument
in a test, where single-threadedness is a property of the harness, and
the synthetic gate asserts it there. It is demonstrably armed rather
than assumed: `test_qwen3_forward.cpp:559` — the assertion itself, not
the `TEST_CASE(` line at `:497` — already asserts on CPU that it reaches
exactly `5 * num_hidden_layers` when the dispatcher does fall back.

**And the counter is not the only process-wide static on this route, so
this description should say what the row does NOT decline.**
`dense_nvfp4::DenseMarlinWorkspace` (`dense_nvfp4_gemm.h:506`) is a
process-static DEVICE allocation — `static void* ws` behind a `static
std::mutex`, sized from `MarlinDeviceSms` and keyed on nothing, not even
the device index — shared by every caller of `MatmulNvfp4MarlinD`
(`:539`) and `GateUpFusedMarlinD` (`:700`). It has three consumers today
(the shared dense route in `dense_attn_block.h`, MiniMax-H3, and the
compressed-tensors NVFP4 scheme); NemotronH's `lm_head` becomes a
fourth. It is PRE-EXISTING — it arrived at `80d1da096` and its
definition is byte-identical at this row's merge base and at its head —
and nothing here addresses it. In a change whose thesis is refusing to
inherit the sibling process-static defect
([#984](#984)), saying nothing
about the one it does inherit was an asymmetry a reviewer is entitled to
see stated. The cost the fallback would carry is 198.18 MB per call —
176.16 MB of packed codes plus 22.02 MB of group scales — arithmetic
that checks to the byte against the `[131072, 2688]` geometry. The
behavioural red for this class needs CUDA and is PENDING with the rest.

### The Marlin arm builds here, and saying otherwise overstated the
blocker

The previous description said "this branch has not built the Marlin path
at all". `include/vt/cuda/marlin_repack.h` includes only `<cstdint>`,
`<cstddef>` and `<vector>`, so the host side needs no CUDA toolkit. On
this box, which has no `nvcc`:

```
c++ -std=c++20 -I include -I src -isystem third_party -DVT_MARLIN_NVFP4=1 \
    -Wall -Wextra -Werror -c -o nhd_marlin.o \
    src/vllm/model_executor/models/nemotron_h_device.cpp
-> rc 0, 0 errors, 0 warnings, a 1 268 552-byte object
```

**That byte count is anchored to the merge commit on this branch**, gcc
13.3.0, no `nvcc`, and compilation is deterministic here — two runs of
the identical command, identical `sha256 42d670b6...`. The anchor is the
correction. The figure this description carried until now, 1 269 696,
was accurate when written and accurate at the reviewed head `29b1128e3`,
but the translation unit moved twice afterwards — 1 272 808 at
`fedf78d86`, 1 268 552 from `bff2b7b2f` on — and an evidence block
naming no SHA cannot tell a reader which of the three it means. The
measured series across seven commits is in the spec.

The size is incidental either way; what the block asserts is `rc 0`,
zero errors and zero warnings on a toolchain with no CUDA, and that
holds at every one of those commits. Both arms are compiled and both are
compiled `-Werror`. What genuinely needs `nvcc` is the Marlin **kernel**
and every **execution** of the device path; those stay PENDING and are
what the table below records.

### Reachability, and a CPU vehicle that already existed

Both synthetic cases build `NemotronHHostWeights` by hand, which proves
the class works and not that anything reaches it, and the deletion
mutation for the device arm needs CUDA. But this row also refactored the
part of `NemotronHPagedForward` that every CPU step runs, and
`tests/vllm/models/test_nemotron_h_paged_forward.cpp` already drives
that function through a real `GPUModelRunner`. It gains a case: three
gathered rows against one request, so the returned row count must follow
`logits_indices` and never `num_reqs`; the trace's copy of
`final_normed` fed back through the production host entry point must
reproduce the returned logits bit for bit; and the same logits must come
back with the trace off.

With `trace->capture` set, `DownloadF32(final_normed)` ran twice on the
host arm. One download now serves both consumers on every path. That
duplicate is not observable from outside the function — both copies are
of the same unchanged buffer — so it is repaired structurally and no
assertion pretends to catch it.

### Records

`spec §4.3` did not exist; the requirement is under `## 5. Owed`. The
allowlist's `## Gates` did not either; the heading is `## 3. The gate`.
`docs/FEATURES.md` still listed the device `lm_head` as owed while
`docs/USAGE.md` said it was device; both now say implemented and never
executed, which is the honest state. The appended issue-index row for
[#1410](#1410) cited `## Owed`
where the heading is `## 5. Owed` — cosmetic, since the row passes on
its owning-row disjunct, but the index is append-only and this is the
only window to correct it.

[#1421](#1421) is filed rather
than fixed in flow: two NVFP4 W4A16 upstream anchors are wrong at the
pin, verified against a checkout of `vllm-project/vllm` at
`5559679229bc961848b121ccdeaa8fa5d79bec98`.
`kernels/linear/__init__.py:879-881` is the tail of
`init_wfp8_a16_linear_kernel`, an FP8 helper sharing the `a16`
vocabulary, and the forced-Marlin line it quotes is at 922-924;
`qwen3.py:271-274` is the `class Qwen3ForCausalLM(...)` line and
`packed_modules_mapping` is at 275-278. Both predate the merge base, and
both are cited from nine code and test sites over SEVEN files plus eight
record sites over FIVE record files (`parity-ledger.md` carries three).
The index row said "six files" while its own parenthetical enumerated
seven, and "six record files" where there are five;
`.agents/issue-index.md` is append-only and a landed row can never be
edited, so an unmerged pull request is the only window to correct that,
and this one does. The set still needs re-deriving at the pin rather
than a sed. This branch relocates one of those citations and
deliberately leaves its VALUE unchanged, so the issue owns every site
uniformly instead of half of them looking reviewed.

## Gate status, stated as pending rather than as passed

Every number here is from a run on a CPU box with no CUDA toolkit,
`RelWithDebInfo`, `-Wall -Wextra -Werror`.

**Read the `overlay` column first.** This tree cannot construct a
`GPUModelRunner` at all —
[#1371](#1371) throws there —
so every green below was taken with
[#1392](#1392 production fix
applied to the working tree, never committed here and reverted
byte-for-byte afterwards. The previous version of this table said "a run
on this tree" and reported the §12 case at `run_rc 0 / SUCCESS!`, which
is false of the tree it named: un-overlaid, that case is already red.

| binary | overlay | `run_rc` | cases | assertions | verdict |
|---|---|---|---|---|---|
| `test_nemotron_h_moe_device` | none | 0 | 4, 4 passed | 4 | `SUCCESS!`
— and **all four SKIP**, both A2-Q2b cases included |
| `test_nemotron_h_paged_forward`, whole binary | **none** | 1 | 13, 2
passed, 11 failed | 18 | `FAILURE!` —
[#1371](#1371), not this row |
| `test_nemotron_h_paged_forward`, whole binary | **#1392** | 0 | 13,
**13 passed** | 3269 | `SUCCESS!` |
| the new §12 case alone | **none** | **1** | 1, **0 passed, 1 failed**,
12 skipped | **0** | **`FAILURE!` — throws #1371 before its first
assertion** |
| the new §12 case alone | **#1392** | 0 | 1, 1 passed, 12 skipped | 13
| `SUCCESS!` |

**The `moe_device` row is not a pass, and it is the honest state of the
synthetic gate.** The binary builds and exits 0, but every case takes
the `TryCudaQueue` skip on a GPU-less box, so those 4 assertions are the
skip notices themselves. The numeric gate examined nothing.

**The paged-forward red is
[#1371](#1371 All 11
failures throw the identical `No valid attention backend for device type
0 from {FLASH_ATTN: [head_size not supported]}` at `GPUModelRunner`
construction, and 10 of them are cases this row never touched.
Overlaying [#1392](#1392
production fix in the working tree — never committed here, reverted
afterwards — turns the same binary green. Worth recording: doctest
printed `assertions: 18 | 18 passed | 0 failed` while 11 cases were
throwing, so the assertion line alone would have read as a pass.

### The red-first, on the cases that can run

Each mutation went into a scratch copy of `nemotron_h_device.cpp`, was
proven applied by `git diff --stat`, was built before being run (a
mutation that fails to build proves nothing), and was restored to an
identical sha256.

**Every row below was also run with
[#1392](#1392) overlaid**, and
that is not a detail. Un-overlaid, the §12 case is already `run_rc=1` on
the UNMUTATED tree, so an M1 or M2 red measured there would prove
nothing whatever. The overlay is what lets the vehicle report a green in
the first place; only then can a mutation take it away.

| mutation | overlay | applied | `compile_rc` | `run_rc` | verdict |
|---|---|---|---|---|---|
| — (unmutated control) | **#1392** | — | 0 | 0 | **GREEN**, `SUCCESS!`,
13 assertions |
| — (unmutated control) | **none** | — | 0 | **1** | already **RED** —
#1371, which is why the rest is overlaid |
| **M1** — `n_out` -> `R` at the host projection | **#1392** | 1 ins / 1
del | 0 | **1** | **RED**, `FAILURE!` |
| **M2** — never fill `trace->final_normed` | **#1392** | 1 ins / 2 del
| 0 | **1** | **RED**, `FAILURE!` |
| **M4** — restore the exact pre-repair two-download shape | **#1392** |
3 ins / 8 del | 0 | 0 | **GREEN — reported, not hidden** |

Re-derived rather than restated, on a clean tree at the merge commit,
overlay and mutations reverted to byte-identical sources
(`nemotron_h_device.cpp` back to `sha256 abf6e21f...`): §12 alone
un-overlaid gives `run_rc=1`, `1 | 0 passed | 1 failed | 12 skipped`,
`assertions: 0`, `THREW exception: No valid attention backend for device
type 0 from {FLASH_ATTN: [head_size not supported]}`; overlaid,
`run_rc=0`, 1 passed, 13 assertions; the whole binary overlaid, 13/13,
3269 assertions; M1 `compile_rc=0 run_rc=1` throwing `gathered row count
does not match hidden_size` at `nemotron_h.cpp:1029`; M2 `compile_rc=0
run_rc=1` on `REQUIRE( 0 == 288 )`. M4's cell needs no separate
attestation — un-overlaid every run of this case is red, so a GREEN is
only reachable with the overlay and its verdict entails the column.

M1 is the red the `n_out` rename exists for: with the request count
substituted, the returned row count is 1 where the gather asked for 3.
M2 arms the trace-operand assertion. Both showed the same trap — doctest
printed `assertions: 2 | 2 passed | 0 failed` on M1 while the case was
failing, because a `REQUIRE` throws rather than counting.

M4 staying green is a result, not a gap. The duplicate `DownloadF32`
copies the same unchanged buffer twice and produces identical bytes, so
nothing observable from outside the function can distinguish it. It is
repaired structurally and no assertion pretends to catch it.

**The device arm's own red-first does not exist on a CPU box and is not
claimed.** `MarlinW4A16Selects` is false on a CPU queue, so the device
branch, the `fallback_gemms` assertion and the reachability deletion
mutation are all unreachable here.

### A gate this row was already failing, and cannot repair in place

`scripts/check-doc-checkpoint.py` is red on this branch, and it was red
at the reviewed head `29b1128e` with the identical two errors (`--base
96ed834 --head 29b1128`):

```
ERROR: commit 1c62d99: changed user_usage but did not update docs/USAGE.md
ERROR: commit 8fa900a: changed .agents/benchmark-record.md: measurement
       recorded but did not update docs/STATUS.md
```

The previous description said this gate had one real failure and that it
was repaired. It was not, and `.github/workflows/ci.yml:519` runs the
same `--base/--head` invocation, so the lane is red for this reason
independently of
[#1371](#1371).

The second error names a real gap and it is now closed: this row moves a
lifecycle state and records a measurement, and `docs/STATUS.md` said
neither. It does now, in 192 characters, because that page carries a
shrink-only ratchet on oversized cells and the first attempt took it
from 44 to 45.

The first error, and the historical form of the second, cannot be closed
by a later commit. The checker iterates `commits_in_range` and judges
each commit on its own contents, so the obligation belongs to
`1c62d9974` and `8fa900a62`. Discharging it means rewriting commits that
are already the reviewed base, which resets this pull request's CI
approval and moves the head a fresh reviewer was asked to look at.
**That is a scheduling decision, so it is recorded as owed and raised
for the operator rather than taken by a repair pass.** Run over this
repair's own commits (`origin/main..HEAD`), the checker reports these
two and nothing else.

### The spec sent the next GPU window through the wrong mutex

`§4` step 1 read "Take `$GPU_LOCK` with a blocking `flock` and wait".
`AGENTS.md` requires a fleet device to be claimed through `rc`, and
`dgx:gpu0` is one. The text predates that rule, but this branch is the
wrong place to leave it standing: it ADDS
`scripts/nemotron-h-a2q2b-gpu-gate.sh` for exactly that window, the
script correctly takes no mutex of its own, and the contradicting
instruction sat beside it in the same file — and every PENDING leg above
is waiting on that window, so the next reader of `§4` is the person
about to open it. Two mutexes that do not exclude each other are worse
than one: the fleet cannot see `$GPU_LOCK`, so a `flock` over `ssh` does
not exclude a concurrent `rc` holder and the controller reports the box
free while somebody is on it. On 2026-08-17 that pair voided a whole
speed axis (`minimax-music3.md` §13.10). Step 1 now claims through `rc
run -d dgx:gpu0 --max-runtime <N>h --` with the gate script as the
payload, and step 2's headroom check moves inside the lease and names
the script's own PRECONDITION 1.

### The trailer gates had stopped running

At the previous head, `origin/main` was not an ancestor, so
`scripts/agent-preflight.sh --fail-on-skip` took its `TRAILER_BEHIND`
arm and SKIPPED both `commit-trailers` and `commit-style`. A skip is not
a green. **It then happened five more times during the device-leakage
repair**, because `origin/main` moved under it at `9ecaf1bb3`,
`f07f96e1c`, `aeba0de6f`, `01854663c` and `c8d926ea8`. The ref only
moves on a fetch, so the fix is ordering: merge, then do not fetch
again, then gate, then push. At the pushed head both gates RUN:
`BASE_SHA=c8d926ea8` non-empty, `ANCESTRY_STATUS=0`, `RANGE_COUNT=22`,
`RANGE_STATUS=0`, **zero `SKIP` tokens anywhere in the report**, 85
gates `ok`, and the only failure is `doc-checkpoint range` with the two
commits above. The skip arms are falsified individually rather than read
off a green line, because `agent-preflight.sh` prints a summary only on
failure. Each is armed over exactly this range, proven by a detached
scratch commit with a period-terminated subject and no trailer block,
which turns `check-commit-style.py` and `check-commit-trailers.py` red
with the specific messages.

Both files that both sides of the merge touched are records, resolved by
shape rather than by whatever the three-way merge produced.
`.agents/issue-index.md` is a genuine append-only log carrying
`merge=union`, and both sides append at the tail, so all four new rows
survive — verified additive, 469 + 2 lines on each side to 471, no line
removed or altered. `docs/USAGE.md` is a keyed record: the merged file
is byte-identical to `origin/main`'s version with this branch's scoped
edit re-applied, so every key neither side owns is unchanged.

### The DSR ratchet was red, and no review round ever saw it run

`device-leakage` never COMPLETED while this pull request was under
review, so its verdict was an input to none of the three passes. It
completed after the third and failed: `vt_ifdef` **35 against a baseline
of 32**, `rc 1`. Three `#ifdef VT_MARLIN_NVFP4` sites had been added to
the device-agnostic shared layer, which is exactly what the ratchet
exists to stop.

They were located by running `scripts/check-device-leakage.py --report`
at the failing head `7a3909187` and diffing the per-file table against
`9ecaf1bb3`, rather than by reading the diff for guards.

| # | site at `7a3909187` | what the guard decided | resolution |
|---|---|---|---|
| 1 | `include/vllm/model_executor/models/dense_nvfp4_gemm.h:768` —
inside `MarlinW4A16Selects` | **nothing** | **removed** |
| 2 | `src/vllm/model_executor/models/nemotron_h_device.cpp:883` —
around `LmHeadNvfp4View` | **nothing** | **removed** |
| 3 | `src/vllm/model_executor/models/nemotron_h_device.cpp:985` —
`DeviceLmHeadD`'s body | `ResidentIn` and a complete
`dense_nvfp4::MarlinDenseResident` | **`DSR-ALLOW(A2-Q2b)`** |

**No baseline was changed.** `AGENTS.md` forbids making a red gate green
by widening an assertion, and a baseline bump is that.
`scripts/device-leakage-baseline.json` is untouched at 32, and the
per-file table is now byte-identical to `origin/main`'s.

**(1) is the case the checker's own message describes.**
`MarlinW4A16Selects` is a SELECTION wearing a build guard, and every
term it reads exists in every build: `MarlinW4A16Enabled()` is declared
above the guarded region, and `vt::OpRegistered` is the op/provider
table's own answer to whether the Marlin arm is realized for a device.
The flag and the registration are one condition, not two —
`CMakeLists.txt`'s single `if(VLLM_CPP_MARLIN)` block adds
`src/vt/cuda/cuda_moe_marlin.cu`, whose file-scope `Registrar` holds the
tree's only `RegisterOp(OpId::kMoeGroupedGemmNvfp4Marlin, …)`, and
defines `VT_MARLIN_NVFP4=1` in that same block. A build without the
macro registers nothing, so the query already resolves false on exactly
the builds the `#ifdef` excluded. This is the call
`nemotron_h_device.cpp`'s `moe_on_device` selection had already made, in
a comment that says so.

**(2) was measured rather than reasoned.** `LmHeadNvfp4View` names
nothing the Marlin build adds — `Nvfp4Weight` comes from
`qwen3_5_weights.h` and `OwnedBytes` from the loader, both
unconditional, and no `vt::cuda::` symbol appears in it. The claim that
its external linkage at `namespace vllm` scope is what makes an unused
definition harmless where its only call site is compiled out is proven
by the mutation that removes that property: adding `static` turns the
same CPU compile RED at `rc 1`, `error: 'vllm::Nvfp4Weight
vllm::LmHeadNvfp4View(...)' defined but not used
[-Werror=unused-function]`. The file was restored to an identical sha256
(`9719ea70…`) afterwards.

**(3) is TYPES-not-behaviour and takes the checker's documented escape
hatch.** `DSR-ALLOW` is not a baseline change: the site is excluded from
the count but COUNTED AND PRINTED on every run, so the exemption is
visible in CI output rather than invisible in the diff.
`DeviceLmHeadD`'s body names two symbols that do not EXIST without the
guarded arena region — the `ResidentIn` template, defined inside it, and
`dense_nvfp4::MarlinDenseResident`, which the header declares
unconditionally and defines only under `VT_MARLIN_NVFP4`, so the
reference cannot bind to an incomplete type. It is the same class and
the same stated reason as the five sibling guards A2-Q2a and A2-P
already carry in this file, and its `#else` refuses by name. The
SELECTION for this arm stays a runtime op-table query.

**Measured on this tree**, a CPU build with `VT_MARLIN_NVFP4` absent
from `build/compile_commands.json` — positive control: 1020 `VLLM_CPP`
hits in the same file, so the grep is not silently wrong — which is the
configuration that exercises both removals, because it is the arm the
deleted `#else` branches used to serve.

| what | before (`7a3909187`) | after | `rc` |
|---|---|---|---|
| `check-device-leakage.py` `vt_ifdef` | 35 | **32** | 1 → **0** |
| `DSR-ALLOW` exemptions in force | 20 | **21** | — |
| `scripts/device-leakage-baseline.json` | 32 | **32, untouched** | — |
| per-file table vs `origin/main` | +1 header, +2 model TU |
**identical** | — |
| `nemotron_h_device.cpp`, `nemotron_h.cpp`, `qwen3_5.cpp` at `-Wall
-Wextra -Werror` | — | compile | **0** |

Each repair is individually load-bearing, proven by reverting it alone
in a scratch worktree and re-running the gate. Every mutation was
verified applied by `git diff --stat` and restored to an identical
sha256, with the unmutated control green immediately before and after.

| mutation | applied | `vt_ifdef` | `rc` | verdict |
|---|---|---|---|---|
| — (control) | — | 32 | 0 | `ratchet holds` |
| **M-B** — restore the guard on `MarlinW4A16Selects` | 4 ins | **33** |
**1** | **RED**, `DSR REGRESSION` |
| **M-C** — restore the guard around `LmHeadNvfp4View` | 2 ins | **33**
| **1** | **RED**, `DSR REGRESSION` |
| **M-D** — delete the `DSR-ALLOW(A2-Q2b)` line | 1 del | **33** | **1**
| **RED**, `DSR REGRESSION` |
| — (control, after restore) | — | 32 | 0 | `ratchet holds` |

The two pre-existing allowlist entries the report also prints,
`deepseek_v4_device.cpp [kcuda] x8` and `platform.cpp [dev_cast] x1`,
are byte-identical at `9ecaf1bb3` and here. This change moves one bucket
and nothing else.

**`docs/USAGE.md` rides in the same commit, because it has to and
because it is true.** Any commit touching `include/vllm/` is
`user_usage` to `scripts/check-doc-checkpoint.py` and owes the surface
in that same commit. `doc-checkpoint range` is already red on this
branch for `1c62d9974` and `8fa900a62`, which an operator decision
covers; a third error of the identical kind would be new damage rather
than inherited, so the header hunk was recommitted with the USAGE edit
beside it. The edit is not written to feed the gate: the NemotronH arms
table enumerates what selects the host `lm_head` projection and omitted
`VT_NVFP4_MARLIN=0`, a user-settable knob that selects an arm and
appears nowhere else in USAGE, and this change is precisely what makes
that knob reach the model's eligibility test in every build rather than
only where the guard compiled it in. `check-doc-checkpoint.py --commit`
is `OK` on that commit, with the still-red `1c62d9974` as the positive
control that the checker is armed.

**`origin/main` moved four times during this repair** — `9ecaf1bb3`,
`f07f96e1c`, `aeba0de6f`, `01854663c` — and each time `git merge-base
--is-ancestor` answered 1, which is the arm that SKIPS both trailer
gates without failing. Each was merged rather than worked around, and
every auto-merged keyed record was resolved by key on every merge: the
branch delta across the merge byte-identical to the branch delta from
the merge base, checked file by file with a positive control on the
comparison itself.

### The three red CPU jobs were inherited, and the merge cleared them

`build-test-cpu` and both `sanitize-cpu` arms failed on this pull
request, all three on `test_nemotron_h_paged_forward` —
[#1371](#1371), which
[#1392](#1392) fixed on `main`
after the review rounds. `sanitize-cpu` is NOT a no-baseline job, so
that had to be triaged by WHICH test failed rather than assumed.

The prediction is verified rather than asserted. With `origin/main`
merged, the same binary is green on this tree with NO overlay: `13 | 13
passed | 0 failed`, `assertions: 3269 | 3269 passed | 0 failed`,
`Status: SUCCESS!`, `run_rc 0` — the same numbers the spec's tables
previously recorded only WITH #1392 applied to the working tree and
reverted. **The `#1392` overlay is therefore retired**; the historical
rows in the spec keep their `overlay` cells, because they describe the
tree they were measured on and rewriting them would make them false.

### `main` brought four red suites with it, and they are attributed both
ways

The third sync carried `4712dac40` (`VT-ACT-ROUND-POLARITY`,
[#1322](#1322) via
[#1347](#1347)), and with it four
red CPU suites this branch does not touch. Both `ctest` runs are on this
box, no overlay of any kind.

| tree | base | `ctest` | result |
|---|---|---|---|
| this row + `9ecaf1bb3` | before `4712dac40` | 567 | **`100% tests
passed, 0 tests failed out of 567`**, `rc 0` |
| this row + `01854663c` | after `4712dac40` | 569 | `rc 8`, **4
failed**: `test_minimax_music3_ar`, `test_ltx2_text_encoder`,
`test_muse_glimmer_text`, `test_muse_glimmer_text_fallback` |

Those four are [#1458](#1458),
filed from another flow before this control was run. Inherited, not
caused, and proven in both directions on the same tree rather than
argued — each mutation verified applied by `git diff --stat` and
restored to an identical sha256:

| control | change | `compile_rc` | `run_rc` (music3 / ltx2 / glimmer) |
verdict |
|---|---|---|---|---|
| **A** | revert THIS row's two source files to their pre-repair
`7a3909187` content | 0 | **1 / 1 / 1** | still RED — **not this row's**
|
| **B** | revert `src/vt/cpu/cpu_ops.cpp` alone to `4712dac40^` | 0 |
**0 / 0 / 0**, `37/37`, `27/27`, `24/24` | GREEN — **`4712dac40` is the
cause** |

Control A answers the attribution question on its own; control B is here
because naming a cause is more useful to the next reader than clearing
oneself. Neither repairs anything, deliberately: #1458 needs
`VT-ACT-ROUND-POLARITY` to decide whether its kernel or four bf16 error
floors that were never re-derived are the wrong side, which is that
row's oracle work rather than a small and clear in-flow fix.

**That window has since closed.**
[#1458](#1458) was fixed on
`main` and closed, and this branch has merged the repair, so the four
are not expected on the current head. The controls stay recorded because
a red that was attributed and then disappeared is still the reason two
`ctest` runs in this description disagree, and deleting the evidence
would leave that unexplained.

### `documentation-checkpoint` is the same two commits, and it is not
new

It went red for the same reason `device-leakage` did: it had never
COMPLETED during review. Its whole failure set is `1c62d9974` and
`8fa900a62`, the two an operator decision covers and which cannot be
repaired in place — `check-doc-checkpoint.py` iterates
`commits_in_range` and judges each commit on its own contents, so a
later commit cannot close either, and discharging them means rewriting
the reviewed base.

The job's three invocations were reproduced locally at this head with
CI's own arguments (`--base` the PR base `63d87805c`, `--head` HEAD):

| invocation | result |
|---|---|
| `check-doc-checkpoint.py` | `rc 1` — **exactly `1c62d9974` and
`8fa900a62`, nothing else** |
| `check-now-current.py` | `rc 0`, `OK: .agents/NOW.md is a current,
in-budget resume digest` |
| `check-role-discipline.py` | `rc 0`, `OK: every change on main arrived
on a task branch` |

**This repair's own commits are clean, and that is not an accident.**
The first draft of the header fix added a THIRD error of the identical
kind, and it is the reason `docs/USAGE.md` rides in the same commit as
the `include/vllm/` hunk rather than in a follow-up: the checker's
obligation is per-commit, so a follow-up cannot discharge it.
`check-doc-checkpoint.py --commit 65ab066` is `OK`, with the still-red
`1c62d9974` as the positive control that the checker is armed against
exactly that class.

### Nothing ran at all for a whole push, and the cause was a conflict
rather than a queue

The push before this one produced **zero check runs** on `2c63cc87a`
while other pull requests were starting normally. `gh pr view` named it:
`mergeable=CONFLICTING`, `mergeStateStatus=DIRTY`. GitHub could not
build `refs/pull/1415/merge`, so no `pull_request` workflow had anything
to run against. A healthy Actions queue plus a branch with zero runs is
the shape a conflict makes, and reading it as a slow queue would have
cost the whole verdict.

The conflict was `docs/STATUS.md` and `docs/USAGE.md`, and neither was
resolved by picking a side, because both had been **relocated** on
`main`. `USAGE.md` is now a 231-line hub whose per-model content moved
under `docs/models/`, and `STATUS.md` is a 27-row surface summary with
no per-model rows at all. So both were taken WHOLE from `main` and this
row's scoped edits re-applied BY KEY at their new home,
`docs/models/nemotron-3-5-lightning.md` — 22 insertions and exactly ONE
deletion there, every unrelated key byte-identical.

One of this row's edits is **dropped rather than re-applied**,
deliberately, and the other was never dropped at all — the claim that
both were is corrected below. The `STATUS.md` Nemotron row is genuinely
gone: merged `main`'s `STATUS.md` is a 98-line surface summary that
enumerates no model at all. `Nemotron` returns 0 hits there, and so do
the positive controls `Kimi` and `Laguna`, while the file plainly greps.
Re-adding it would re-create a row a relocation removed, which is the
duplicate a second relocation makes. The `USAGE.md`
`NemotronHForCausalLM` refusal row is the opposite case, and the earlier
claim that it "no longer exists anywhere under `docs/`" was **wrong**:
`aee6c48d6` (#1491) did not delete that row, it **RELOCATED** it to
`docs/reference/model-loading.md` — a file that did not exist at this
branch's pre-merge base, which is precisely why the earlier pass read it
as gone. It is still on `main`, and this row now updates it.
`docs/FEATURES.md` auto-merged and kept this row's cell — checked, not
assumed. `check-conflict-markers.py` (#1450, new on `main`) reports 0
findings over 3820 tracked text files.

Run locally from the ORIGINAL merge base, `check-doc-checkpoint.py` also
reports `d995c52f0`, `af25bd251` and `995ed1ccd` — each verified an
ancestor of `origin/main`, none in this branch's own range, all arrived
with the merge. CI does not see those, because it resolves `--base` to
the CURRENT merge base: **its whole failure set is `1c62d9974` and
`8fa900a62`**, the two an operator decision covers, and nothing else.
This row's own commits stay clean either way: `check-doc-checkpoint.py
--commit` is `OK` on `65ab06636` and on the merge itself.

### Two public rows still said the head runs on the host, and this
change is what made them false

Following the relocation above to its consequence:
`docs/reference/model-loading.md:188` still read "`lm_head` and FP8
Mamba2 projections run on the host". That sentence was true before this
row and is false after it, which is the test for whose obligation it is.
Left alone, the tree would land with
`docs/models/nemotron-3-5-lightning.md` calling the head a device arm
and this row calling it a host arm, on the same merge commit.

Only the clause this row falsified moves. `lm_head` reaches the device
**on the paged forward** and is carried as implemented-and-unmeasured,
matching what the model recipe and the spec's `## Now` already say. The
FP8 Mamba2 projections are a **different unit** (A2-Q1, #1289, still
held DRAFT) and still run on the host, so that clause is kept and only
re-worded to stop sharing a verb with the head. The pending-token-gate
and no-GGUF clauses are untouched, because both are still accurate.

Sweeping by CLAIM rather than by file, at the MERGED tree, found a
**second** site the scoped review did not reach, because it is not a
refusal row. The checkpoint registry at `docs/USAGE.md:230` lists
per-checkpoint arms, and its "Supported arms" cell read "host FP8 Mamba2
and NVFP4 head". That "NVFP4 head" is this same `lm_head`, so the same
test applies and the same one clause moves. That cell is the surface the
"say which weights, and from where" rule owns, so a supported-arms cell
putting the head on the host is exactly the contradiction this finding
is about.

The sweep carried its own controls, because a null grep only proves the
terms wrong. Every `docs/**.md` and `README.md` line asserting a host
placement for a head or for logits is 11 hits, 2 of them the sites
repaired here (positive control); the same pipeline with a nonsense term
returns 0 (negative control). The other 9 were read rather than
pattern-matched, and none is stale: `BENCHMARKS.md:16` is this row's own
host-re-expansion ATTRIBUTION and is what motivates the device arm;
`BENCHMARKS.md:276` is Laguna GB10 ATS weight residency;
`ENVIRONMENT.md:144,187,212` are Qwen3.5/Qwen3.6 levers; `server.md:88`
is the CPU-only `prompt_logprobs` full-logits route on the shared
runner, not this model; `model-loading.md:144` is the FP8 scalar-scale
guard; and two are historical `docs/superpowers/plans/` documents, which
are not projections.

**No gate can catch either repair.** `check-doc-checkpoint.py` asks
whether SOME `docs/USAGE.md` edit accompanied a commit of the changed
class. It never asks whether a sentence inside it is still true.

### Every check has COMPLETED on the current head, and the remaining
reds are all inherited

`ci` run `32537115264` on `99f9f672a`: **25 checks, 0 pending** — 14
pass, 6 fail, 5 skipped by design on a pull request (`attest`,
`baseline-summary`, `manifest`, `promote`, `publish`).

| check | on `b7d89b43e` | on `99f9f672a` | whose |
|---|---|---|---|
| **`device-leakage`** | SUCCESS | **SUCCESS** | **this row's.** The
restructure survives the 38-commit merge at `vt_ifdef` 32, and
`scripts/device-leakage-baseline.json` is byte-identical to `main`, so
the row adds no new allowance |
| `build-newest-gcc` | FAILURE | **SUCCESS** | was **`main`'s**, and it
cleared exactly as predicted once `main` carrying #1581 was merged. CI
builds the MERGE commit, so the red came from the base side; #1565 is
closed |
| `build-test-cpu`, `sanitize-cpu (address,undefined)`, `sanitize-cpu
(thread)` | SUCCESS | **FAILURE** | **`main`'s**,
[#1608](#1608) and
[#1602](#1602) — see below |
| `build-test-vulkan`, `build-test-cpu-arm64`, `cuda-fat-build`,
`cuda-arch-features`, `vulkan-spirv-freshness`, `agent-record`,
`pr-size`, `commit-protocol-tag`, `last-gated-commit`, `plan`, `verify`
x2 | SUCCESS | SUCCESS | — |
| `documentation-checkpoint` | FAILURE | FAILURE | `1c62d9974` +
`8fa900a62` ONLY — the two an operator decision covers, verified against
the job log rather than assumed |
| `windows-msvc-cpu`, `windows-msvc-vulkan` | FAILURE | FAILURE | the
two no-baseline PR-only jobs,
[#584](#584) |

The three CPU legs all fail on **`282 - test_runner`**, and the failure
is inherited rather than this row's. The failing case, `runner:
initialize_kv_cache refuses a non-multiple-of-16 block size`, has **0**
occurrences at `b7d89b43e` — the head where all three legs were SUCCESS
— and **1** at the merge base `5453e571d`. It was introduced by
`e2a9e035d` (#1273), which **is** among the 38 commits the merge brought
in, and this row's range touches **0** files in that test (control: it
does touch the two `tests/vllm/models/test_nemotron_h_*.cpp` files, so
the query discriminates).

It is also **not a data race**, which matters because `sanitize-cpu
(thread)` is the lane that would expose one. All three legs report the
SAME deterministic assertion — `test_runner.cpp:1557`, `20 test cases /
19 passed / 1 failed`, `544 assertions / 543 passed / 1 failed` — and
**0** sanitizer findings in any of them (`WARNING: ThreadSanitizer` = 0,
`ERROR: AddressSanitizer` / `runtime error:` = 0). A race cannot make a
non-sanitized `build-test-cpu` fail identically. The assertion expects
`"Block size must be a multiple of 16"` but the attention registry
refuses first, with `"No valid attention backend for device type 0 from
{CPU_ATTN: [block_size not supported], FLASH_ATTN: [block_size not
supported]}"` — which is precisely what #1602 describes. Nothing is
suppressed here and nothing is pragma'd, because there is nothing of
this row's to suppress.

### Gate table

| leg | state |
|---|---|
| seam extension, device arm, production wiring | DONE; built on BOTH
arms, `-Werror`, 0 warnings |
| the CPU-reachable half, gated through a real `GPUModelRunner` | DONE
and RUN, with M1/M2 red-first |
| routing allowlist narrowed,
[#1410](#1410) filed | DONE |
| synthetic device `lm_head` numeric gate | COMPILES; CUDA-only; **NEVER
RUN** |
| `nvcc` build of the Marlin kernel, real-checkpoint numeric leg, token
identity, reachability deletion mutation, the `VT_NVFP4_MARLIN=0`
fallback red | PENDING a `dgx:gpu0` window, job queued |

A reviewer should treat every CUDA leg as unrun. The spec's `## 6. Now`
carries the same tables, so the pending state lives with the row rather
than only in this description.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]

---------

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>

… the correctness-grade attention kernel (#1545) (#1579)

`src/vllm/model_executor/models/muse_glimmer_vision.cpp:639` called
`vt::Attention`
for every segment of every one of the perception encoder's 50 blocks.
`src/vt/ops.cpp:2680`
resolves that op to the kernel whose own header at
`src/vt/cuda/cuda_ops.cu:1456-1460`
calls itself "Correctness-grade (M0.9)": one 256-thread block per
(query, head), a
256-wide shared-memory tree reduction for every key, and no K/V tiling.
The tower has
no second attention path, no knob and no A/B rung, so nothing could
route around it. It
now names `vt::AttentionDenseFlash`, the rung
`whisper_audio.cpp:310-322` and
`qwen3_vl_vision.cpp:462-480` already default to. `AttentionDenseFa2` is
not usable at
head_dim 96 and would fall through to the same kernel anyway.

The size of the problem is now read rather than inferred. #1545 labelled
its token count
illustrative and asked whoever took it to pin the shipped resolution
first, because it
moves the estimate quadratically, and the released checkpoint was
already on the NAS.
`config.json` ships `layer_types` explicitly -- 13 `full_attention` and
37
`window_attention`, exactly what `muse_glimmer_weights.cpp:429-439`
derives when the key
is absent -- and `pos_emb_height/width` put the window at 32x32 = 1024
patches, so a
window layer does not attend over the whole image.
`processor_config.json` sets
`max_image_tokens` to 4096. Whether that counts patch tokens or
post-merge tokens is not
decidable from the checkpoint, so both readings are carried: 34 s or 375
s of naive
attention per image, against the 4.8 s the issue illustrated with.

This lands UNREACHED, deliberately, under the staged-slice exception in
`.agents/reachability.md`. What is not reached:
`MuseGlimmerVisionForward` and its two
callers `MuseGlimmerEncodePixelGroups` and
`MuseGlimmerGenerateGreedyViaRegistry` have no
caller in `src/`, in `examples/` or in `include/vllm.h`, only in
`tests/`, and
`muse_glimmer_registry.cpp:13-14` says why: "The perception encoder is
still W3, so an
image or video prompt is a pending brick." The row that owns the wiring
is
`MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation`, at W4
and W5 of
`.agents/specs/muse-glimmer.md` §3. The issue that tracks it is #1566,
filed by this
change because the model's umbrella #268 is closed and nothing open
named the gap; both
specs list it under `## Owed`. Repairing the cost before the wiring
costs one line.
Repairing it afterwards costs a regression hunt through a change that
did not touch
attention.

Numerics, both halves. On CPU the swap is byte-identical BY
CONSTRUCTION:
`src/vt/cpu/cpu_ops.cpp:3760` registers `kAttentionDenseFlash` to the
same
`AttentionKernel` function pointer `kAttention` gets at `:3750`, and
each op has exactly
one CPU provider. Every existing golden passes at its previous value and
no tolerance was
widened. On CUDA it is NOT bit-identical: the flash kernel reduces
per-lane head_dim
strips with `__shfl_xor_sync` under a running online softmax where the
naive kernel uses a
256-wide block tree, so the f32 partial-sum grouping differs, as
`include/vt/ops.h:3304-3327`
states. No CUDA run backs this change -- `dgx:gpu0` is held by the
operator and no lease
was taken -- and the A/B is recorded as owed rather than claimed.

The test had to be a ROUTING test, because a numeric one cannot see this
change: on CPU
both ops are the same kernel, so any comparison passes before it and
after it. The new
case counts op-provider selections the way
`tests/vllm/models/test_ltx2.cpp:665-686` does
for the same question, and asserts a DERIVED 12 -- 5 window + 2 full + 5
window calls over
the fixture's segmentation -- so a partial re-route fails it too.

`src/vt/cuda/cuda_ops.cu` is deliberately untouched. #1544 item 2 owes a
repair to
`LaunchAttentionDenseFlash`'s advertised `head_dim <= 256` contract and
the LTX-2.5 row is
implementing it. head_dim 96 does not need it: the production bf16 arm
asks 24 KiB of the
48 KiB default cap, and the spec records that the f32 arm sits exactly
on the cap and
fails loud rather than silently if it ever runs on CUDA.

## How to verify

```sh
cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release
cmake --build build --target test_muse_glimmer_vision test_muse_glimmer_wiring -j 8
./build/tests/test_muse_glimmer_vision
./build/tests/test_muse_glimmer_wiring
```

Observed on this branch: 8/8 cases and 103 assertions, then 9/9 cases
and 10,317
assertions, both `SUCCESS!`. The routing case prints `attention
selections:
dense-flash +12, naive +0`.

Red first, with the test in place and the production line still on
`vt::Attention`:
`dense-flash +0, naive +12`, and both `CHECK`s fail as `CHECK( 0 == 12
)` and
`CHECK( 12 == 0 )`.

Byte-identity measured, not only argued. Both binaries were built from
one tree
with only that line differing, and every `MESSAGE:` line of the suite
was captured
from each and diffed. Exactly one line of fourteen differs, and it is
the routing
counter. The f32 stage numbers stay at `rel_l2` 1.201e-07 to 2.983e-07
and the
bf16 tower at `rel_l2=5.951e-03 max_abs=3.675e-02` on both sides.

`scripts/agent-preflight.sh` reports **All gates green.** on the head,
run here.
An earlier run of it on this branch failed
`test_cpu_x86_llamacpp_floor`, and that
failure was inherited rather than caused here: it failed the same way on
the
unmodified base `04f1cead6`, with `waiting for quiet: 15s busy=111%
builders=0
load=67.47` and exit 4 where the case wants 2, which is the
load-dependent
behaviour #618 tracks. The host was at loadavg 67 to 121 then and lower
at the
final run, and the gate passed.

THREE CI lanes are red on this pull request and none is caused here.
Each was
verified per-job against a main baseline, by comparing the failing
ASSERTION and
not the job name:

- `windows-msvc-cpu` and `windows-msvc-vulkan` fail at `ERROR:
CMakeLists.txt:
  MSVC /W4 /WX policy is negated on the C/C++ compile by /w`, before any
compilation. They are red on main's own scheduled baseline too, and #584
/ #965
  track them.
- `build-test-cpu` and both `sanitize-cpu` arms fail on ONE shared
doctest case,
`test_runner.cpp:1557`, which landed on main in `e2a9e035d` (#1273) and
is owned
by #1602 and #1608. The scheduled main baseline at `e2a9e035d` fails the
identical `CHECK_THROWS_WITH_AS( make_runner(), "Block size must be a
multiple
of 16", std::invalid_argument )` with the same "No valid attention
backend for
device type 0" text, in all three jobs, with zero sanitizer findings in
either
  sanitized arm. This branch does not touch `test_runner.cpp`.

`build-newest-gcc` was red when this body was first written and is GREEN
here:
#1581 and #1618 landed the `::getpid` repair on main and this branch was
re-merged onto `2e7f3bee7` to pick it up. `test_cpu_threadpool`, which
an earlier
run could not verify as inherited, PASSED on the re-run, leaving
`test_runner`
as the whole of the `build-test-cpu` red — 1 failing test out of 585.

## What is not verified

No CUDA run. `dgx:gpu0` was held by the operator, no lease was taken,
and the A/B
that would confirm the speed is recorded under `## Owed` in the spec
rather than
claimed. The cost figures above are arithmetic from the 5.70
ns/iteration GB10
constant recorded in #1544, not a measurement of this tower.

Closes #1545.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]

---------

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>

…de kernel, and one forward cost 47.84 s (#1557)

One LTX-2.5 DiT forward at `768x448/49f` measured **47.84 s** on GB10 --
n=119, median 47.91 s, spread 5.8%, from the engine's own `last=`
samples -- and first-order arithmetic says it should take well under a
second.

The DiT self-attention called `vt::Attention`, which on CUDA resolves to
the kernel whose own header calls itself "Correctness-grade (M0.9)": one
256-thread block per (query, head), a 256-wide shared-memory tree
reduction per key, and no K/V tiling, so K and V are re-read from global
once per (query, head). At that geometry the video stream is 2352 tokens
x 32 heads = 75,264 blocks each looping 2352 keys, over 48 layers.

The attribution is arithmetic rather than assertion.
`.agents/specs/multimodal-speed.md` measures that same kernel on that
same box at **5.70 ns per block-key iteration**, and 1.77e8 x 5.70 ns x
48 = **48.4 s** against the measured 47.84 s. A 1% match, which leaves
the other 1% for 48 blocks of GEMMs, norms, RoPE, gating and six
attentions each.

**After this change one forward is 7.680 s** (n=19, median,
`768x448/49f`, GB10). The ratio is **~6.0x**, and it is **not** an A/B.
Both qualifications are load-bearing and both are under "The A/B" below.

## Why nobody saw it, which is the part worth keeping

`kAttention` is deliberately frozen on the naive kernel so that text
decode stays byte-identical. That decision is correct and this change
does not touch it.

The consequence is the defect. The fast kernels are **separate ops that
each caller must opt into by name** -- `kAttentionDenseFast`,
`kAttentionDenseFlash`, `kAttentionDenseFa2`. There is no automatic
selection, no shape routing, and no fallback notice. A model that never
opts in gets **correct output at roughly 500x the cost, with no warning
anywhere**.

Nothing in this tree can detect that. The output is right, so every
golden passes. The op is registered, so no refusal fires.
`GetOpProviderStats` counts the naive selection and reports it as a
success, because it is one. The only symptom is a wall clock, and a
diffusion render has no reference wall clock to be held to. That is why
the gate for this change is a dispatch observation and not a number.

## What changed

**`ltx2_device.cpp` calls `vt::AttentionDenseFlash` in the
self-attention branch.** That is the whole of the product change. The
dispatch RULE is unchanged: the branch is still chosen by `context ==
nullptr && a.bias == nullptr`, upstream's own self-attention marker,
never by what the numbers happen to be. Only the op it calls moves. Its
square-problem contract holds by construction, because that branch is
entered only when `s == tq`.

**`VLLM_LTX2_DIT_FLASH_ATTN=0` restores the old op**, so both arms of
the measurement run from one binary. Same shape as `VT_FA2_DENSE`,
documented in `docs/ENVIRONMENT.md` as a measurement lane and never a
configuration. Both remaining naive `vt::Attention` call sites -- that
one, and `ltx2.cpp`'s CPU-only host arm -- carry a `// VT-ATTN-NAIVE:`
line saying why, the form #1578's checker defines for a deliberate site.

**The host arm at `ltx2.cpp` is deliberately NOT moved.** It computes
into `std::vector<float>` and is CPU-only by construction, and on CPU
both ops are the same registered function
(`src/vt/cpu/cpu_ops.cpp:3750-3761`), so the swap would be a
byte-identical no-op that moves the L2 parity reference off the
reference op.

## The shared-memory cap-raise is REVERTED, and that is the change since
the review

An earlier version of this branch also raised
`LaunchAttentionDenseFlash`'s dynamic shared-memory cap through
`cudaFuncSetAttribute`, moved `SetDynamicSmemOptIn` onto a shared seam,
and added head_dim contract cases to `tests/vt/test_ops_attention.cpp`.
**All of it is gone.** `src/vt/cuda/cuda_ops.cu`, `cuda_device_caps.h`,
`cuda_arch_tactics.cu`, `cuda_paged_attn.cu` and
`tests/vt/test_ops_attention.cpp` are byte-identical to `main`.

**The swap never needed it.** LTX renders at the stream dtype, and in
production that is bf16 (`ltx2_device.cpp:1183`). The flash op's K/V
tile is `2 * kFlashBc(64) * head_dim * sizeof(Tin)`, so the video
stream's head_dim 128 asks **32,768 B** and the audio stream's head_dim
64 asks 16,384 B. Both are inside the **49,152 B** every CUDA
architecture gives a launch without any opt-in at all.

**What the raise was actually serving does not fit.** The shapes over 48
KiB here are f32: head_dim 128 at 65,536 B, and head_dim 256 at 131,072
B against GB10's queried ceiling of **101,376 B**
(`cuda_device_caps.h:46`). The 256 shape therefore does not fit *even
with the opt-in* -- it was falling back to the bit-identical
`AttentionDenseFast` while the case reported a launch. So the raise
bought this row one shape it does not run, and paid for it by moving a
shared helper across two files and colliding with #1578 on the same
lines.

**#1578 owns the bound and takes the opposite, better approach.** Rather
than raising the cap it makes the ADVERTISED domain honest:
`AttentionDenseFlash` declares `head_dim <= 256` while it can only
launch bf16 192 / f32 96, and #1578 narrows the declaration to what the
code can do and refuses above it. That is a property of the code rather
than of whichever device is underneath, and it is the
`supports_head_size()` polarity vLLM already has. **#1578 merges
first**, and after it bf16 head_dim 128 is inside the declared bound, so
this change is unaffected. There is now nothing left to conflict.

**One consequence is disclosed rather than left to be found.** With the
raise gone, the f32 L2 parity arm at production geometry (head_dim 128,
65,536 B) reaches `AttentionDenseFlash` and cannot launch: a
`cudaGetLastError` throw at `cuda_ops.cu:3352` today, a `VT_CHECK`
naming the head_dim once #1578 lands. **It fails loud in both worlds and
never silently**, and nothing gated reaches it -- production is bf16,
and the f32 arm is a parity reference exercised at the fixture's reduced
dimensions. Filed under `## Owed`.

## Numerics, measured rather than asserted

**On CPU: byte-identical.** `kAttention` and `kAttentionDenseFlash` are
the same registered function pointer, and the goldens are unmoved --
`test_ltx2_device` **22/22, 652/652**; `test_ltx2` **43/43, 4581/4581**;
`test_ltx2_video` **102/102, 4194/4194**; `test_ops_attention_cross`
**9/9, 32/32**.

**On CUDA: NOT bit-identical, and here is the number.** The warp kernel
groups the head_dim partial sums across 32 lanes instead of a 256-thread
block, so the same f32 online softmax associates differently. On
`dgx:gpu0` the host-vs-device parity case measures **video 8.9407e-08,
audio 4.47035e-08** against its committed `2e-5`, and bf16
CUDA-vs-CPU-backend at **0**. That is f32 round-off scale, 224x inside
the gate, and **the gate was not widened**.

**That is the only numeric evidence there is, and it bounds less than it
looks like it bounds.** The case runs the fixture's reduced dimensions,
so it bounds the ARITHMETIC change -- a length-D sum reassociated -- and
not the change at head_dim 128 with 2352 keys over 48 layers. A
diffusion render has no token gate to fall back on, and the flash arm
was interrupted before writing any frames, so **no pixel comparison
exists** either, not even against the completed 49-frame baseline render
already on the NAS. Filed as **#1612** and listed under `## Owed`.

**The `test_ops_attention` evidence is WITHDRAWN, not restated.** The
GB10 lease ran it at 10/10 and 88,439 assertions, and this change claims
nothing from that run: it measured the head_dim cases that came with the
cap-raise, and those are no longer in the tree. Two of its three arms
would not have supported the claim anyway, for the reason above. The
`test_ltx2_device` rows and the render below DO survive the revert, and
that is arithmetic rather than assertion -- the measured binary carried
the opt-in call, but the helper returns immediately below 49,152 B, so
it was a **no-op on every launch those numbers came from**.

## Reachability, twice

**Unit.** A new case drives the production entry point
`Ltx2DitForwardDevice` -- called from the denoise loop at
`ltx2_video.cpp:4246` -- and asserts the dispatch **two-sidedly**
through `GetOpProviderStats`: `kAttentionDenseFlash` selected exactly
**8** times (two self-attentions x two blocks x two batch rows) and
`kAttention` selected **0**. The negative half is what makes it a
routing proof rather than an addition proof. The case scope-guards its
own process state, because it enables a counting instrument and sets an
env var and contains `REQUIRE`s; measured with a scratch
`REQUIRE(false)` and an appended observer, **without the guards the
observer reads the knob still set and 8 leaked selections, and with them
it reads neither**.

That `:4246` is itself a repair. This branch and its spec both cited
`ltx2_video.cpp:4055-4059`, which points at prose about the res_2s step
counter -- at the row's declared base `6b48edb2c` and at HEAD alike, so
re-reading it at either revision would have caught it. It is the
citation the whole reachability argument rests on.
`cpu_ops.cpp:3551-3562` was wrong the same way and pointed at
`FusedStore`. Every other anchor in the spec has been re-read against
the base rather than carried forward, and eight more were corrected.

**Mutation M1**: restore `vt::Attention` at the call site, +1/-1,
compile rc 0. Both halves went red -- `CHECK( 0 == 8 )` and `CHECK( 8 ==
0 )`, exit 1 -- while every golden case in the same binary stayed green.
That contrast is the finding: no numerical gate in this tree can see a
500x slower kernel that computes the right answer. Tree restored and
re-gated green.

**Production, on the real model.** With `VT_OP_PROVIDER_STATS=1` the
full 21.00B render at `768x448/49f` on GB10 announces `op=21 device=1`
(`kAttentionDenseFlash` on CUDA) and announces `op=18 device=1`
(`kAttention` on CUDA) **zero times**. Same two-sided claim, taken
through `--device cuda` at full scale rather than on a fixture.

## The A/B: one arm measured, the pair still PENDING

Lease `6c724dfd` on `dgx:gpu0`, source `30dce3a1d`, one binary built
in-lease with cutlass-nvfp4, cutlass-fp8 and FA-2 all `ENABLED for
[121a]`. Correctness cleared before any speed number was read.

**Flash arm**, per DiT forward at `768x448/49f` = 2352 tokens, from the
engine's own `last=` lines:

| n | median | mean | min | max | spread |
|---|---|---|---|---|---|
| 19 | **7.680 s** | 7.633 s | 7.109 s | 8.196 s | 14.2% |

**The naive arm did not run, so this is not an A/B.** At forward 20 the
`rc` worker was lost and `dgx:gpu0` read `unhealthy (no contact)`. The
cause is UNPROVEN and this change does not name one: no memory trace was
taken and the box did not return to be asked. What is established is
that the harness as first written carried no memory guard and no sample
cap, which is a defect in this row's harness rather than a finding about
the change.

**The 47.84 s denominator carries two confounds this arm does not, and
both inflate the ratio.** Neither was disclosed before:

- **A stack sampler.** The denominator ran under `runguard.py
--stack-period 12` (`render.log:1`), which `eu-stack`s the process and
so `ptrace`-stops every thread. Its own `stacks.txt` prices that: **523
samples, median inter-sample delta 12.40 s against a 12.0 s period**, so
~0.40 s median and 1.50 s max of stopped process per sample. About 3.9
samples land inside each 47.84 s forward, ~1.54 s, **~3.2%**. Correcting
only the denominator gives **46.3 s / 7.680 s = 6.03x**.
- **A different prompt.** `render.log:1` carries a ~70-word prompt; the
harness uses one short sentence, and `ltx2_video.cpp:2253` sets
`context_tokens = encoded.seq` **unpadded**, so the DiT's
cross-attentions see a different number of keys in each arm.
Corroborated rather than inferred: `conditioning.tower` is **45.013 s**
against **28.426 s**. Same sign, and not quantified.

So the defensible statement is the range **6.03x to 6.23x**, quoted as
**~6.0x**, with the sampler correction named and the prompt confound
uncorrected and pushing the same way. `6.23x` survives in the records
only as the uncorrected upper end of that range, never on its own. The
A/B gate reads `PENDING` and this change does not claim otherwise.

**The flash arm's artifacts do not record what it ran**, which is why
those confounds had to be established from a phase duration.
`arm-flash.log` opens at `[render] + load` with no command line,
`wd-flash/` is empty, no `phase-log.json` was written, and the only
description of the run was a mutable NAS path edited 25 minutes after it
finished. Both halves are repaired: the harness is committed as
**`scripts/ltx25-dit-attn-flash-ab.sh`**, and every arm now writes its
own invocation -- harness sha256, binary sha256, source SHA, geometry,
seed, prompt, resolved command line -- to line 1 of its own log. It also
caps each arm at 13 samples, holds a 12 GiB `MemAvailable` floor, caches
the build on the source SHA, and runs the **naive arm first**.

That harness also had a precondition that could never pass. It grepped
`cuda_ops.cu` for `FlashTileSmemOptIn`, a spelling no revision of this
change ever used, so it counted 0 and would have `exit 42`-ed on a
correct tree as readily as on a wrong one. It is removed with the
cap-raise it guarded.

## CI

- **`documentation-checkpoint` was RED and it was THIS BRANCH's, not
inherited.** `2aa78c69b` and `2f39a9426` each recorded a measurement in
`.agents/benchmark-record.md` without writing `docs/STATUS.md` (and
`docs/BENCHMARKS.md` for the second); the control on the main-only range
`4c193bd55..5d548d0` is rc 0. Neither commit is in this branch's
history any more, and the checker is **re-run at this head** rather than
trusted to have stayed fixed -- a job that has stopped appearing in a
failing set is not the same fact as a job that passes. Local rc 0 over
the branch range. A side effect worth recording: that job runs `set -eu`
and this checker is the FIRST of three commands, so
**`check-now-current.py` and `check-role-discipline.py` never ran in CI
on this branch at all**. Both are rc 0 locally, so nothing hides behind
it.
- **`build-newest-gcc` is now GREEN on `main`** since #1581 landed, and
this branch carries that fix through the merge. Earlier runs of this PR
predate it. A red here now would be this branch's, not inherited.
- `build-test-cpu` and both `sanitize-cpu` lanes: **inherited**, and
verified from this head's own logs rather than from the issue numbers.
All three fail on exactly one case out of 585 -- `test_runner.cpp:1557`,
`CHECK_THROWS_WITH_AS(make_runner(), "Block size must be a multiple of
16", ...)` receiving `No valid attention backend for device type 0`
instead -- with byte-identical text in all three, and `main`'s own
newest baseline run fails the same case. `test_ltx2_device` passes in
all three, so this change's own cases are green on the lanes that run
them. From #1273; owned by #1602 and #1608.
- `windows-msvc-cpu` / `windows-msvc-vulkan`: **inherited**,
baseline-less lane. A markdown-only control PR (#1295) fails the
identical step. #584/#965 own them.
- The full set: **16 pass, 5 fail**, and the five are a strict SUBSET of
`main`'s newest baseline at `503e459005d7`. `scripts/main-baseline.py`
was the instrument, not the push runs, which are all cancelled (#274).

## Owed, filed and not folded in

- **#1612** -- there is no numeric or pixel comparison at production
geometry. The only numeric gate is the reduced-dimension one above; the
flash arm wrote no frames, so no pixel A/B exists. The f32 parity arm's
refusal at production geometry is recorded against the same issue.
- **#1551** -- `vt::AttentionDenseFa2` still refuses head_dim 128, so
LTX cannot reach tensor cores. Everything here is still a scalar
warp-per-query recurrence.
- **#1552** -- the same opt-in-by-name defect reaches every other
`vt::Attention` caller, and nobody will be told there either.

All three are listed under `## Owed` in
`.agents/specs/ltx25-dit-attn-flash.md`, together with the two
`scripts/attention-rung-allowlist.txt` stems that #1578's checker will
report `STALE` once the markers here meet it.

Closes #1549

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]

---------

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>

…ority this row waited on

Merging `main` at `8eecc05a9` invalidated three claims in this row's own
records. They are repaired here rather than left for a reader, because a record
that contradicts its tree is the failure this protocol exists to prevent.

FIRST, the inherited-red row is now green. `build-test-cpu` and both
`sanitize-cpu` lanes were recorded as INHERITED, failing on one case at
`test_runner.cpp:1557` from #1273. PR #1700 fixed #1608 on `main` by declaring
ROCM_ATTN's MultipleOf(16) block sizes, and the merge carried it: all three
report `pass` on head `3f31bf0d4`, and `main` itself now reports `success` on
`build-test-cpu` and `sanitize-cpu (thread)`. An inherited-red claim that has
gone green does not merely age; it would have excused a future red on three
lanes that are clean.

SECOND, the flash arm is no longer the default. #1551 landed as `a4f2a9585`
while this row was in review, instantiating FA-2 at head_dim 128 and routing the
DiT onto its tensor cores, so `VLLM_LTX2_DIT_FLASH_ATTN` is now three-way with
FA-2 unset-default and `=flash` the arm section 10.7 measured. Nothing in 10.7
is invalidated - it is a same-binary naive-against-flash pair at a recorded
binary sha256, and it is the denominator #1551's own ratio is taken against.
What moves is the reach of the QUESTION: section 10.4 wrote V4 as
self-calibrating precisely so "the FA-2 arm (#1551) can take this same
criterion", and #1743 now asks it of a rung this row never rendered. A pixel A/B
of the FA-2 arm is owed and is not this row's.

THIRD, the complete CI rollup is recorded at its exact SHA: 18 pass, 6 skip, 3
fail on `3f31bf0d4`, every failure characterised BY TEXT. Both Windows lanes
print `Windows portability contract OK` and then die at
STATUS_STACK_BUFFER_OVERRUN in test_openai_api_server.exe after the length-40
refusal, which is #584 and not #1649 - #1701 rescoped that gate. `agent-record`
prints `agent record OK: ENGINE=170 ...` and then FileNotFoundError: 'hugo',
which is #1722. The twelve lanes green here that are not inherited are named, so
a future red on any of them is attributable rather than arguable.

The merge authority is recorded as a gate rather than assumed. It stood PENDING
while an agent relayed "merge it", because a relayed instruction is not
developer consent, and the refusal cost time and was correct. The developer gave
it directly on 2026-08-23. `.agents/developer-preferences.md` carries it scoped
to this row and explicitly not as blanket authority.

The `## Owed` entry for FA-2 head_dim 128 takes main's HANDED OFF wording; the
two bullets main does not carry are preserved. Every unrelated key in
`docs/ENVIRONMENT.md` is byte-for-byte main's, with only the one knob row
re-edited.

The row stays ACTIVE, not DONE, while #1743 is open.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]